#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 1;
class KernelInstanceNorm
{
public:
  __aicore__ inline KernelInstanceNorm() {}
  __aicore__ inline void Init(GM_ADDR x, GM_ADDR gamma, GM_ADDR beta, GM_ADDR y, GM_ADDR mean, GM_ADDR variance, uint32_t blockLength,
                              uint32_t tileNum, uint32_t tileLength, uint32_t lasttileLength, 
                              uint32_t typeKey, uint32_t totalLength, uint32_t iValue, uint32_t jValue, uint32_t valueCount, float epsilon,
                              uint32_t gammaLength, uint32_t betaLength)
  {
    ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
    this->typeKey = typeKey;
    this->minLength = 32 / sizeof(DTYPE_X);
    this->valueCount = 
        valueCount ASSERT(valueCount != 0 && "data_format doesn't match the data shape");
    this->iValue = iValue;
    this->jValue = jValue;
    this->gammaLength = gammaLength;
    this->betaLength = betaLength;
    this->totalLength = totalLength;

    this->blockLength = blockLength;
    this->tileNum =
        tileNum ASSERT(tileNum != 0 && "tile num can not be zero!");
    this->tileLength = tileLength / BUFFER_NUM;
    this->lasttileLength = lasttileLength;
    this->x_addr = x;
    this->gamma_addr = gamma;
    this->beta_addr = beta;

    if(typeKey == 0){
      xGm.SetGlobalBuffer((__gm__ float *)x ,
                        this->iValue * this->jValue * this->valueCount);
      gammaGm.SetGlobalBuffer((__gm__ float *)gamma,
                        this->iValue * this->jValue * this->valueCount);
      betaGm.SetGlobalBuffer((__gm__ float *)beta,
                        this->iValue * this->jValue * this->valueCount);
      yGm.SetGlobalBuffer((__gm__ float *)y,
                        this->iValue * this->jValue);
      meanGm.SetGlobalBuffer((__gm__ float *)mean,
                        this->iValue * this->jValue);
      varianceGm.SetGlobalBuffer((__gm__ float *)variance,
                        this->iValue * this->jValue);

      pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(inQueueG, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(inQueueB, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(inQueueT, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(inQueueT2, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(outQueueM, BUFFER_NUM, this->minLength * sizeof(float));
      pipe.InitBuffer(outQueueV, BUFFER_NUM, this->minLength * sizeof(float));
      pipe.InitBuffer(calcBuf, BUFFER_NUM * this->minLength * sizeof(float));
    }
    else{
      xGm_h.SetGlobalBuffer((__gm__ half *)x ,
                        this->iValue * this->jValue * this->valueCount);
      gammaGm_h.SetGlobalBuffer((__gm__ half *)gamma,
                        this->iValue * this->jValue * this->valueCount);
      betaGm_h.SetGlobalBuffer((__gm__ half *)beta,
                        this->iValue * this->jValue * this->valueCount);
      yGm_h.SetGlobalBuffer((__gm__ half *)y,
                        this->iValue * this->jValue);
      meanGm_h.SetGlobalBuffer((__gm__ half *)mean,
                        this->iValue * this->jValue * 2);
      varianceGm_h.SetGlobalBuffer((__gm__ half *)variance,
                        this->iValue * this->jValue * 2);

      xGm.SetGlobalBuffer((__gm__ float *)x ,
                        this->iValue * this->jValue * this->valueCount);
      gammaGm.SetGlobalBuffer((__gm__ float *)gamma,
                        this->iValue * this->jValue * this->valueCount);
      betaGm.SetGlobalBuffer((__gm__ float *)beta,
                        this->iValue * this->jValue * this->valueCount);
      yGm.SetGlobalBuffer((__gm__ float *)y,
                        this->iValue * this->jValue);
      meanGm.SetGlobalBuffer((__gm__ float *)mean,
                        this->iValue * this->jValue);
      varianceGm.SetGlobalBuffer((__gm__ float *)variance,
                        this->iValue * this->jValue);

      pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(inQueueG, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(inQueueB, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(inQueueT, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(inQueueT2, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(float));
      pipe.InitBuffer(outQueueM, BUFFER_NUM, this->minLength * sizeof(float));
      pipe.InitBuffer(outQueueV, BUFFER_NUM, this->minLength * sizeof(float));
      pipe.InitBuffer(calcBuf, BUFFER_NUM * this->minLength * sizeof(float));

      pipe.InitBuffer(inQueueX_H, BUFFER_NUM, this->tileLength * sizeof(half));
      pipe.InitBuffer(inQueueG_H, BUFFER_NUM, this->tileLength * sizeof(half));
      pipe.InitBuffer(inQueueB_H, BUFFER_NUM, this->tileLength * sizeof(half));
      pipe.InitBuffer(inQueueT_H, BUFFER_NUM, this->tileLength * sizeof(half));
      pipe.InitBuffer(inQueueT2_H, BUFFER_NUM, this->tileLength * sizeof(half));
      pipe.InitBuffer(outQueueY_H, BUFFER_NUM, this->tileLength * sizeof(half));
      pipe.InitBuffer(outQueueM_H, BUFFER_NUM, this->minLength * sizeof(half));
      pipe.InitBuffer(outQueueV_H, BUFFER_NUM, this->minLength * sizeof(half));
      pipe.InitBuffer(calcBuf_H, BUFFER_NUM * this->minLength * sizeof(half));
    }
    
  }
  __aicore__ inline void Process()
  {
    if(typeKey == 1){
      ProcessFp16();
    }else{
      ProcessFp();
    }
  }

private:
  __aicore__ inline void ProcessFp16()
  {
    int32_t loopCount = this->tileNum * BUFFER_NUM;
    int32_t i =0, j = 0;
    for(int32_t i = 0; i < iValue; i++)
      for(int32_t j = 0; j < jValue; j++){
        // 一组元素-----得到一个mean值
        for (int32_t k = 0; k < loopCount; k++)// ReduceSum(x)
        {
          CopyInFp16(i, j, k, 0); 
          Compute(i, j, k, 0);
          CopyOutFp16(i, j, k, 0);
        }
        // 前面其实是对一组元素进行求和，最后需要求mean
        // ========CopyIn======
        LocalTensor<float> tmpLocal = inQueueT.AllocTensor<float>();
        DataCopy(tmpLocal, meanGm[this->iValue * this->jValue / 2 + 1], this->minLength);
        inQueueT.EnQue(tmpLocal);
        // ========Compute======
        LocalTensor<float> meanLocal = outQueueM.AllocTensor<float>();
        tmpLocal = inQueueT.DeQue<float>();
        meanLocal.SetValue(0, (float)((float)tmpLocal.GetValue(0) / valueCount));
        inQueueT.FreeTensor(tmpLocal);
        outQueueM.EnQue<float>(meanLocal);
        // ========CopyOut======
        meanLocal = outQueueM.DeQue<float>();
        DataCopy(meanGm[this->iValue * this->jValue / 2 + 1], meanLocal, this->minLength);
        outQueueM.FreeTensor(meanLocal);

        // 一组元素-----得到一个variance值
        for (int32_t k = 0; k < loopCount; k++) // ReduceSum((x_tensor - mean_scalar) * (x_tensor - mean_scalar))
        {
          CopyInFp16(i, j, k, 1);
          Compute(i, j, k, 1);
          CopyOutFp16(i, j, k, 1);
        }
        // variance与mean都得到后，处理得到y-------gamma*((x-mean)/sqrt(vaerance+epsilon))+beta
        // ========CopyIn======
        tmpLocal = inQueueT.AllocTensor<float>();
        DataCopy(tmpLocal, varianceGm[this->iValue * this->jValue / 2 + 1], this->minLength);
        inQueueT.EnQue(tmpLocal);
        // ========Compute======
        LocalTensor<float> varianceLocal = outQueueV.AllocTensor<float>();
        tmpLocal = inQueueT.DeQue<float>();
        varianceLocal.SetValue(0, (float)((float)tmpLocal.GetValue(0) / valueCount));
        inQueueT.FreeTensor(tmpLocal);
        outQueueV.EnQue<float>(varianceLocal);
        // ========CopyOut======
        varianceLocal = outQueueV.DeQue<float>();
        DataCopy(varianceGm[this->iValue * this->jValue / 2 + 1], varianceLocal, this->minLength);
        outQueueV.FreeTensor(varianceLocal);

        for (int32_t k = 0; k < loopCount; k++)// ReduceSum(x)
        {
          CopyInFp16(i, j, k, 2);
          Compute(i, j, k, 2);
          CopyOutFp16(i, j, k, 2);
        }

        tmpLocal = inQueueT.AllocTensor<float>();
        DataCopy(tmpLocal, meanGm[this->iValue * this->jValue / 2 + 1], this->minLength);
        inQueueT.EnQue(tmpLocal);
        // ========Compute======
        LocalTensor<half> meanLocal_h = outQueueM_H.AllocTensor<half>();
        tmpLocal = inQueueT.DeQue<float>();
        meanLocal_h.SetValue(0, (half)(tmpLocal.GetValue(0)));
        inQueueT.FreeTensor(tmpLocal);
        outQueueM_H.EnQue<half>(meanLocal_h);
        // ========CopyOut======
        meanLocal_h = outQueueM_H.DeQue<half>();
        DataCopy(meanGm_h[i*jValue+j], meanLocal_h, this->minLength);
        outQueueM_H.FreeTensor(meanLocal_h);

        tmpLocal = inQueueT.AllocTensor<float>();
        DataCopy(tmpLocal, varianceGm[this->iValue * this->jValue / 2 + 1], this->minLength);
        inQueueT.EnQue(tmpLocal);
        // ========Compute======
        LocalTensor<half> varianceLocal_h = outQueueV_H.AllocTensor<half>();
        tmpLocal = inQueueT.DeQue<float>();
        varianceLocal_h.SetValue(0, (half)(tmpLocal.GetValue(0)));
        inQueueT.FreeTensor(tmpLocal);
        outQueueV_H.EnQue<half>(varianceLocal_h);
        // ========CopyOut======
        varianceLocal_h = outQueueV_H.DeQue<half>();
        DataCopy(varianceGm_h[i*jValue+j], varianceLocal_h, this->minLength);
        outQueueV_H.FreeTensor(varianceLocal_h);
      }
  }

  __aicore__ inline void CopyInFp16(int32_t i, int32_t j, int32_t progress, int32_t kind)
  { 
    if(kind == 0){
      LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
      LocalTensor<float> tmpLocal = inQueueT.AllocTensor<float>();
      LocalTensor<half> tmpLocal_h = inQueueT_H.AllocTensor<half>();
      uint32_t dataLength = tileLength;
      if(progress == tileNum - 1){
        dataLength = lasttileLength;
      }
      uint32_t start_index = i*valueCount*jValue + j + progress * tileLength * jValue;
      for(int p=0; p < dataLength; p++){
        xLocal.SetValue(p, ((__gm__ half *)x_addr)[start_index + p*jValue]);
      }
      
      DataCopy(tmpLocal, meanGm[this->iValue * this->jValue / 2 + 1], this->minLength);
      inQueueX.EnQue(xLocal);
      inQueueT.EnQue(tmpLocal);
      inQueueT_H.FreeTensor(tmpLocal_h);
    }
    else if(kind == 1) {
      LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
      LocalTensor<float> meanLocal = inQueueT.AllocTensor<float>();
      LocalTensor<float> tmpLocal = inQueueT2.AllocTensor<float>();
      LocalTensor<half> tmpLocal_h = inQueueT_H.AllocTensor<half>();
      uint32_t dataLength = tileLength;
      if(progress == tileNum - 1){
        dataLength = lasttileLength;
      }
      uint32_t start_index = i*valueCount*jValue + j + progress * tileLength * jValue;
      for(int p=0; p < dataLength; p++){
        xLocal.SetValue(p, ((__gm__ half *)x_addr)[start_index + p*jValue]);
      }
      
      DataCopy(meanLocal, meanGm[this->iValue * this->jValue / 2 + 1], this->minLength);
      DataCopy(tmpLocal, varianceGm[this->iValue * this->jValue / 2 + 1], this->minLength);
      inQueueT_H.FreeTensor(tmpLocal_h);
      inQueueX.EnQue(xLocal);
      inQueueT.EnQue(meanLocal);
      inQueueT2.EnQue(tmpLocal);
    }
    else{
      LocalTensor<half> xLocal = inQueueX_H.AllocTensor<half>();
      LocalTensor<half> gammaLocal = inQueueG_H.AllocTensor<half>();
      LocalTensor<half> betaLocal = inQueueB_H.AllocTensor<half>();
      LocalTensor<float> meanLocal = inQueueT.AllocTensor<float>();
      LocalTensor<float> tmpLocal = inQueueT2.AllocTensor<float>();
      LocalTensor<half> tmpLocal_h = inQueueT_H.AllocTensor<half>();
      uint32_t dataLength = tileLength;
      if(progress == tileNum - 1){
        dataLength = lasttileLength;
      }
      uint32_t start_index = i*valueCount*jValue + j + progress * tileLength * jValue;
      for(int p=0; p < dataLength; p++){
        xLocal.SetValue(p, ((__gm__ half *)x_addr)[start_index + p*jValue]);
      }
      for(int p=0; p < dataLength; p++){
        gammaLocal.SetValue(p, ((__gm__ half *)gamma_addr)[(start_index + p*jValue)%gammaLength]);
      }
      for(int p=0; p < dataLength; p++){
        betaLocal.SetValue(p, ((__gm__ half *)beta_addr)[(start_index + p*jValue)%betaLength]);
      }
      DataCopy(meanLocal, meanGm[this->iValue * this->jValue / 2 + 1], this->minLength);
      DataCopy(tmpLocal, varianceGm[this->iValue * this->jValue / 2 + 1], this->minLength);
      
      inQueueT_H.FreeTensor(tmpLocal_h);
      inQueueX_H.EnQue(xLocal);
      inQueueG_H.EnQue(gammaLocal);
      inQueueB_H.EnQue(betaLocal);
      inQueueT.EnQue(meanLocal);
      inQueueT2.EnQue(tmpLocal);
    }
  }

  __aicore__ inline void CopyOutFp16(int32_t i, int32_t j, int32_t progress, int32_t kind)
  {
    if(kind == 0){
      LocalTensor<float> meanLocal = outQueueM.DeQue<float>();
      DataCopy(meanGm[this->iValue * this->jValue / 2 + 1], meanLocal, this->minLength);
      outQueueM.FreeTensor(meanLocal);
    }
    else if(kind == 1){
      LocalTensor<float> varianceLocal = outQueueV.DeQue<float>();
      DataCopy(varianceGm[this->iValue * this->jValue / 2 + 1], varianceLocal, this->minLength);
      outQueueV.FreeTensor(varianceLocal);
    }
    else{
      LocalTensor<half> yLocal = outQueueY_H.DeQue<half>();
      LocalTensor<half> tmpLocal_h = inQueueT_H.AllocTensor<half>();
      uint32_t dataLength = tileLength;
      if(progress == tileNum - 1){
        dataLength = lasttileLength;
      }
      uint32_t start_index = i*valueCount*jValue + j + progress * tileLength * jValue;
      uint32_t end_index = i*valueCount*jValue + j + valueCount * jValue;
      if(end_index > start_index + tileLength * jValue){
        end_index = start_index + tileLength * jValue;
      }
      uint32_t now_index = start_index;
      uint32_t index = 0;
      while(now_index < end_index){
        DataCopy(tmpLocal_h, yGm_h[now_index], this->tileLength);
        for(int p=0; p < this->tileLength && now_index + p < end_index && index < dataLength; p+=jValue){
          tmpLocal_h.SetValue(p, yLocal.GetValue(index));
          index++;
        }
        DataCopy(yGm_h[now_index], tmpLocal_h, this->tileLength);

        now_index = now_index + this->tileLength;
        if((now_index - start_index )% jValue){
          now_index = now_index - ((now_index - start_index )% jValue) + jValue;
        }
      }
      outQueueY_H.FreeTensor(yLocal);
      inQueueT_H.FreeTensor(tmpLocal_h);
    }
  }

  __aicore__ inline void ProcessFp()
  {
    int32_t loopCount = this->tileNum * BUFFER_NUM;
    for(int32_t i = 0; i < iValue; i++)
      for(int32_t j = 0; j < jValue; j++){
        // 一组元素-----得到一个mean值
        for (int32_t k = 0; k < loopCount; k++)// ReduceSum(x)
        {
          CopyInFp(i, j, k, 0); 
          Compute(i, j, k, 0);
          CopyOutFp(i, j, k, 0);
        }
        // 前面其实是对一组元素进行求和，最后需要求mean
        // ========CopyIn======
        LocalTensor<float> tmpLocal = inQueueT.AllocTensor<float>();
        DataCopy(tmpLocal, meanGm[i*jValue+j], this->minLength);
        inQueueT.EnQue(tmpLocal);
        // ========Compute======
        LocalTensor<float> meanLocal = outQueueM.AllocTensor<float>();
        tmpLocal = inQueueT.DeQue<float>();
        meanLocal.SetValue(0, (float)((float)tmpLocal.GetValue(0) / valueCount));
        inQueueT.FreeTensor(tmpLocal);
        outQueueM.EnQue<float>(meanLocal);
        // ========CopyOut======
        meanLocal = outQueueM.DeQue<float>();
        DataCopy(meanGm[i*jValue+j], meanLocal, this->minLength);
        outQueueM.FreeTensor(meanLocal);

        // 一组元素-----得到一个variance值
        for (int32_t k = 0; k < loopCount; k++) // ReduceSum((x_tensor - mean_scalar) * (x_tensor - mean_scalar))
        {
          CopyInFp(i, j, k, 1);
          Compute(i, j, k, 1);
          CopyOutFp(i, j, k, 1);
        }
        // variance与mean都得到后，处理得到y-------gamma*((x-mean)/sqrt(vaerance+epsilon))+beta
        // ========CopyIn======
        tmpLocal = inQueueT.AllocTensor<float>();
        DataCopy(tmpLocal, varianceGm[i*jValue+j], this->minLength);
        inQueueT.EnQue(tmpLocal);
        // ========Compute======
        LocalTensor<float> varianceLocal = outQueueV.AllocTensor<float>();
        tmpLocal = inQueueT.DeQue<float>();
        varianceLocal.SetValue(0, (float)((float)tmpLocal.GetValue(0) / valueCount));
        inQueueT.FreeTensor(tmpLocal);
        outQueueV.EnQue<float>(varianceLocal);
        // ========CopyOut======
        varianceLocal = outQueueV.DeQue<float>();
        DataCopy(varianceGm[i*jValue+j], varianceLocal, this->minLength);
        outQueueV.FreeTensor(varianceLocal);

        for (int32_t k = 0; k < loopCount; k++) // res=gamma*((x-mean)/sqrt(vaerance+epsilon))+beta
        {
          CopyInFp(i, j, k, 2);
          Compute(i, j, k, 2);
          CopyOutFp(i, j, k, 2);
        }
      }
        
  }

  __aicore__ inline void CopyInFp(int32_t i, int32_t j, int32_t progress, int32_t kind)
  { 
    if(kind == 0){
      LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
      LocalTensor<float> tmpLocal = inQueueT.AllocTensor<float>();
      uint32_t dataLength = tileLength;
      if(progress == tileNum - 1){
        dataLength = lasttileLength;
      }
      uint32_t start_index = i*valueCount*jValue + j + progress * tileLength * jValue;
      
      for(int p=0; p < dataLength; p++){
        xLocal.SetValue(p, ((__gm__ float *)x_addr)[start_index + p*jValue]);
      }
      
      DataCopy(tmpLocal, meanGm[i*jValue+j], this->minLength);
      inQueueX.EnQue(xLocal);
      inQueueT.EnQue(tmpLocal);
    }else if(kind == 1) {
      LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
      LocalTensor<float> meanLocal = inQueueT.AllocTensor<float>();
      LocalTensor<float> tmpLocal = inQueueT2.AllocTensor<float>();
      uint32_t dataLength = tileLength;
      if(progress == tileNum - 1){
        dataLength = lasttileLength;
      }
      uint32_t start_index = i*valueCount*jValue + j + progress * tileLength * jValue;
      for(int p=0; p < dataLength; p++){
        xLocal.SetValue(p, ((__gm__ float *)x_addr)[start_index + p*jValue]);
      }
      
      DataCopy(meanLocal, meanGm[i*jValue+j], this->minLength);
      DataCopy(tmpLocal, varianceGm[i*jValue+j], this->minLength);
      inQueueX.EnQue(xLocal);
      inQueueT.EnQue(meanLocal);
      inQueueT2.EnQue(tmpLocal);
    }else if(kind == 2) {
      LocalTensor<float> xLocal = inQueueX.AllocTensor<float>();
      LocalTensor<float> gammaLocal = inQueueG.AllocTensor<float>();
      LocalTensor<float> betaLocal = inQueueB.AllocTensor<float>();
      LocalTensor<float> meanLocal = inQueueT.AllocTensor<float>();
      LocalTensor<float> tmpLocal = inQueueT2.AllocTensor<float>();
      uint32_t dataLength = tileLength;
      if(progress == tileNum - 1){
        dataLength = lasttileLength;
      }
      uint32_t start_index = i*valueCount*jValue + j + progress * tileLength * jValue;
      for(int p=0; p < dataLength; p++){
        xLocal.SetValue(p, ((__gm__ float *)x_addr)[start_index + p*jValue]);
      }
      for(int p=0; p < dataLength; p++){
        gammaLocal.SetValue(p, ((__gm__ float *)gamma_addr)[(start_index + p*jValue)%gammaLength]);
      }
      for(int p=0; p < dataLength; p++){
        betaLocal.SetValue(p, ((__gm__ float *)beta_addr)[(start_index + p*jValue)%betaLength]);
      }
      
      DataCopy(meanLocal, meanGm[i*jValue+j], this->minLength);
      DataCopy(tmpLocal, varianceGm[i*jValue+j], this->minLength);
      inQueueX.EnQue(xLocal);
      inQueueG.EnQue(gammaLocal);
      inQueueB.EnQue(betaLocal);
      inQueueT.EnQue(meanLocal);
      inQueueT2.EnQue(tmpLocal);
    }
  }
  __aicore__ inline void Compute(int32_t i, int32_t j, int32_t progress, int32_t kind)
  {
    if(kind == 0){
      LocalTensor<float> xLocal = inQueueX.DeQue<float>();
      LocalTensor<float> tmpLocal = inQueueT.DeQue<float>();
      LocalTensor<float> meanLocal = outQueueM.AllocTensor<float>();
      LocalTensor<float> wLocal = calcBuf.Get<float>();
      if(progress == 0)
        Duplicate(tmpLocal, (float) 0.0f, this->minLength);
      if(progress == tileNum - 1){
        for(int iX = this->lasttileLength; iX < this->tileLength; iX++){
          xLocal.SetValue(iX, (float)0);
        }
      }
      ReduceSum(meanLocal, xLocal, wLocal, this->tileLength);
      meanLocal.SetValue(0, (float)((float)meanLocal.GetValue(0) + (float)tmpLocal.GetValue(0)));
      outQueueM.EnQue(meanLocal);
      inQueueX.FreeTensor(xLocal);
      inQueueT.FreeTensor(tmpLocal);
    }else if(kind == 1){
      LocalTensor<float> xLocal = inQueueX.DeQue<float>();
      LocalTensor<float> meanLocal = inQueueT.DeQue<float>();
      LocalTensor<float> tmpLocal = inQueueT2.DeQue<float>();
      LocalTensor<float> varianceLocal = outQueueV.AllocTensor<float>();
      LocalTensor<float> wLocal = calcBuf.Get<float>();
      if(progress == 0)
        Duplicate(tmpLocal, (float) 0.0f, this->minLength);
      
      Adds(xLocal, xLocal, (float)(0.0f-(float)meanLocal.GetValue(0)), this->tileLength);
      Mul(xLocal, xLocal, xLocal, this->tileLength);
      if(progress == tileNum - 1){
        for(int iX = this->lasttileLength; iX < this->tileLength; iX++){
          xLocal.SetValue(iX, (float)0);
        }
      }
      ReduceSum(varianceLocal, xLocal, wLocal, this->tileLength);
      varianceLocal.SetValue(0, (float)((float)varianceLocal.GetValue(0) + (float)tmpLocal.GetValue(0)));
      outQueueV.EnQue(varianceLocal);
      inQueueX.FreeTensor(xLocal);
      inQueueT.FreeTensor(meanLocal);
      inQueueT2.FreeTensor(tmpLocal);
    }else if(typeKey == 0){
      LocalTensor<float> xLocal = inQueueX.DeQue<float>();
      LocalTensor<float> gammaLocal = inQueueG.DeQue<float>();
      LocalTensor<float> betaLocal = inQueueB.DeQue<float>();
      LocalTensor<float> meanLocal = inQueueT.DeQue<float>();
      LocalTensor<float> varianceLocal = inQueueT2.DeQue<float>();
      LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
      LocalTensor<float> wLocal = calcBuf.Get<float>();
      Adds(xLocal, xLocal, (float)(0.0f-(float)meanLocal.GetValue(0)), this->tileLength);
      // Duplicate(wLocal, (float) sqrt(varianceLocal.GetValue(0) + epsilon), this->tileLength);
      // Div(xLocal, xLocal, wLocal, this->tileLength);
      Muls(xLocal, xLocal, (float)(1.0f / sqrt((float)varianceLocal.GetValue(0) + epsilon)), this->tileLength);
      Mul(xLocal, xLocal, gammaLocal, this->tileLength);
      Add(yLocal, xLocal, betaLocal, this->tileLength);
      outQueueY.EnQue(yLocal);
      inQueueX.FreeTensor(xLocal);
      inQueueG.FreeTensor(gammaLocal);
      inQueueB.FreeTensor(betaLocal);
      inQueueT.FreeTensor(meanLocal);
      inQueueT2.FreeTensor(varianceLocal);
    }else{
      LocalTensor<half> xLocal = inQueueX_H.DeQue<half>();
      LocalTensor<half> gammaLocal = inQueueG_H.DeQue<half>();
      LocalTensor<half> betaLocal = inQueueB_H.DeQue<half>();
      LocalTensor<float> meanLocal = inQueueT.DeQue<float>();
      LocalTensor<float> varianceLocal = inQueueT2.DeQue<float>();
      LocalTensor<half> yLocal = outQueueY_H.AllocTensor<half>();
      LocalTensor<half> wLocal = calcBuf_H.Get<half>();
      
      Adds(xLocal, xLocal, (half)(0.0f-(float)meanLocal.GetValue(0)), this->tileLength);
      Duplicate(wLocal, (half) ((float)varianceLocal.GetValue(0) + epsilon), this->tileLength);
      Sqrt(wLocal, wLocal, this->tileLength);
      Div(xLocal, xLocal, wLocal, this->tileLength);
      Mul(yLocal, xLocal, gammaLocal, this->tileLength);
      Add(yLocal, yLocal, betaLocal, this->tileLength);

      outQueueY_H.EnQue(yLocal);
      inQueueX_H.FreeTensor(xLocal);
      inQueueG_H.FreeTensor(gammaLocal);
      inQueueB_H.FreeTensor(betaLocal);
      inQueueT.FreeTensor(meanLocal);
      inQueueT2.FreeTensor(varianceLocal);
    }
    
  }
  __aicore__ inline void CopyOutFp(int32_t i, int32_t j, int32_t progress, int32_t kind)
  {
    if(kind == 0){
      LocalTensor<float> meanLocal = outQueueM.DeQue<float>();
      DataCopy(meanGm[i*jValue+j], meanLocal, this->minLength);
      outQueueM.FreeTensor(meanLocal);
    }else if(kind == 1){
      LocalTensor<float> varianceLocal = outQueueV.DeQue<float>();
      DataCopy(varianceGm[i*jValue+j], varianceLocal, this->minLength);
      outQueueV.FreeTensor(varianceLocal);
    }else{
      LocalTensor<float> yLocal = outQueueY.DeQue<float>();
      LocalTensor<float> tmpLocal = inQueueT.AllocTensor<float>();
      uint32_t dataLength = tileLength;
      if(progress == tileNum - 1){
        dataLength = lasttileLength;
      }
      uint32_t start_index = i*valueCount*jValue + j + progress * tileLength * jValue;
      uint32_t end_index = i*valueCount*jValue + j + valueCount * jValue;
      if(end_index > start_index + tileLength * jValue){
        end_index = start_index + tileLength * jValue;
      }
      uint32_t now_index = start_index;
      uint32_t index = 0;
      while(now_index < end_index){
        uint32_t tmpindex = index;
        DataCopy(tmpLocal, yGm[now_index], this->tileLength);
        for(int p=0; p < this->tileLength && now_index + p < end_index && tmpindex < this->tileLength; p+=jValue){
          tmpLocal.SetValue(p, yLocal.GetValue(tmpindex));
          tmpindex++;
        }
        DataCopy(yGm[now_index], tmpLocal, this->tileLength);
        index = tmpindex;

        now_index = now_index + this->tileLength;
        if((now_index - start_index )% jValue){
          now_index = now_index - ((now_index - start_index )% jValue) + jValue;
        }
      }
      outQueueY.FreeTensor(yLocal);
      inQueueT.FreeTensor(tmpLocal);
    }
  }

private:
  TPipe pipe;
  TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX, inQueueG, inQueueB, inQueueT, inQueueT2;
  TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX_H, inQueueG_H, inQueueB_H, inQueueT_H, inQueueT2_H;
  TBuf<TPosition::VECCALC>calcBuf, calcBuf_H;
  TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY, outQueueM, outQueueV;
  TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueM_H, outQueueV_H, outQueueY_H;
  GlobalTensor<float> xGm;
  GlobalTensor<float> gammaGm;
  GlobalTensor<float> betaGm;
  GlobalTensor<float> yGm;
  GlobalTensor<float> meanGm;
  GlobalTensor<float> varianceGm;
  GlobalTensor<half> xGm_h;
  GlobalTensor<half> gammaGm_h;
  GlobalTensor<half> betaGm_h;
  GlobalTensor<half> yGm_h;
  GlobalTensor<half> meanGm_h;
  GlobalTensor<half> varianceGm_h;
  GM_ADDR x_addr;
  GM_ADDR gamma_addr;
  GM_ADDR beta_addr;

  uint32_t blockLength;
  uint32_t tileNum;
  uint32_t tileLength;
  uint32_t totalLength;
  uint32_t lasttileLength;
  uint32_t typeKey;
  uint32_t iValue;
  uint32_t jValue;
  uint32_t valueCount;
  uint32_t minLength;
  uint32_t gammaLength;
  uint32_t betaLength;
  float epsilon;
};
extern "C" __global__ __aicore__ void instance_norm(GM_ADDR x, GM_ADDR gamma, GM_ADDR beta, GM_ADDR y, GM_ADDR mean, GM_ADDR variance, GM_ADDR workspace, GM_ADDR tiling) {
  GET_TILING_DATA(tiling_data, tiling);
  KernelInstanceNorm op;

  op.Init(x, gamma, beta, y, mean, variance, tiling_data.blockLength,
          tiling_data.tileNum, tiling_data.tileLength,
          tiling_data.lasttileLength, tiling_data.typeKey, tiling_data.totalLength,
          tiling_data.iValue, tiling_data.jValue, tiling_data.valueCount, tiling_data.epsilon,
          tiling_data.gammaLength, tiling_data.betaLength);
  op.Process();
}

#ifndef __CCE_KT_TEST__
void instance_norm_do(uint32_t blockDim, void* l2ctrl, void* stream,
                       uint8_t* x, uint8_t* gamma, uint8_t* beta, uint8_t* y, uint8_t* mean, uint8_t* variance,
                       uint8_t* workspace, uint8_t* tiling) {
  instance_norm<<<blockDim, l2ctrl, stream>>>(x, gamma, beta, y, mean, variance, workspace, tiling);
}
#endif